#include <iostream>

#define LOG_ARGS(fmt, ...)                                                                                             \
    printf("[blockIdx x %d y %d z %d threadIdx x %d y %d z %d]:" fmt, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, \
           threadIdx.y, threadIdx.z, __VA_ARGS__)
#define LOG(fmt)                                                                                                       \
    printf("[blockIdx x %d y %d z %d threadIdx x %d y %d z %d]:" fmt, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, \
           threadIdx.y, threadIdx.z)

template <typename... Args>
__device__ __forceinline__ void cudaLog(const char *fmt = "", Args &&...args) {
    // 打印固定的线程/块索引头
    char full_fmt[1024] = "[B(%d,%d,%d)T(%d,%d,%d)]";
#define STRLEN(s)            \
    ({                       \
        const char *p = (s); \
        size_t len = 0;      \
        while (*p++) len++;  \
        len;                 \
    })
    auto prefix_len = STRLEN(full_fmt);
    auto fmt_len = STRLEN(fmt);
#undef STRLEN
    for (auto i = 0; i < fmt_len; i++) {
        full_fmt[prefix_len + i] = fmt[i];
    }
    if (sizeof...(Args) > 0) {
        printf(full_fmt, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z,
               std::forward<Args>(args)...);
    } else {
        printf(full_fmt, blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z);
    }
}

__device__ void print_matrix(float* data, int rows, int cols) {
    for (int i = 0; i < rows; i++) {
        for (int j = 0; j < cols; j++) {
            if (j == cols - 1) {
                cudaLog("%10f\n", data[i * 4 + j]);
            } else {
                cudaLog("%10f, ", data[i * 4 + j]);
            }
        }
    }
    cudaLog("\n");
}

#define ROWS 4
#define COLS 4
#define DATA_SIZE ROWS *COLS

__global__ void dummy_kernel(float *input_data, int N) {
    int thread_id = blockIdx.x * blockDim.x + threadIdx.x;

    auto data = input_data + thread_id * DATA_SIZE;
    print_matrix(data, ROWS, COLS);
    cudaLog("\n");
}

#define BLOCK_SIZE 4

int main() {
    const size_t grid_size = 4;
    auto block_size = BLOCK_SIZE;

    const size_t N = block_size * grid_size * DATA_SIZE;
    float h_A[N];

    // 生成正弦波测试数据
    for (int i = 0; i < N; ++i) {
        h_A[i] = sin(i);
    }

    // 设备内存分配
    float *d_A;
    cudaMalloc(&d_A, N * sizeof(float));
    cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
    // 启动核函数

    dim3 grid(grid_size, 1, 1);
    dim3 block(block_size, 1, 1);

    dummy_kernel<<<grid, block>>>(d_A, N);

    // 错误检查
    cudaError_t cu_error = cudaDeviceSynchronize();
    if (cu_error != cudaSuccess) {
        std::cerr << "CUDA error: " << cudaGetErrorString(cu_error) << std::endl;
    }
    // 清理资源
    cudaFree(d_A);

    return 0;
}